
#include "fun.h"
#include "hs_cuda_reduce.cuh"
#include <limits>

template<typename T>
struct ReduceSumFunctor
{
    __device__ __forceinline__ T operator()(T a, T b) const
    {
        return a + b;
    }
};

float reduceSum(const hs::raster::RasterData<float>& src_datas)
{
    hscuda::raster::RasterData<float> device_datas;
    hscuda::raster::host2device(src_datas.view(), device_datas);

    int N = device_datas.view().numel();
    float res_sum = 0.0f;

    hs::cuda::reduce::apply(N
        , device_datas.view().pData()
        , ReduceSumFunctor<float>()
        , res_sum    // 初值
        , &res_sum); // 规约操作结果(返回数据在主机端)
    cudaDeviceSynchronize();

    return res_sum;
}

template<typename T>
struct ReduceMaxIndexFunctor
{
    __device__ __forceinline__ void operator()(T a, int aid, T b, int bid, T* output, int* output_index) const
    {
        if (a > b) {
            (*output) = a;
            (*output_index) = aid;
        }
        else {
            (*output) = b;
            (*output_index) = bid;
        }
    }
};

void reduceMaxIndex(const hs::raster::RasterData<float>& src_datas, float * p_max_val, int * p_max_id)
{
    /* 将输入数据拷贝到的GPU */
    hscuda::raster::RasterData<float> device_datas;
    hscuda::raster::host2device(src_datas.view(), device_datas);

    int N = device_datas.view().numel();

    // 设置规约初值
    float init_val = std::numeric_limits<float>::min();

    hs::cuda::reduceIndex::apply(N
        , device_datas.view().pData()
        , ReduceMaxIndexFunctor<float>()
        , init_val
        , p_max_val
        , p_max_id);
    cudaDeviceSynchronize();
}

